#include <cassert>
#include <chrono>
#include <cstdio>

#define CHECK(cmd)                                                             \
    {                                                                          \
        cudaError_t error = cmd;                                               \
        if (error != cudaSuccess)                                              \
            printf("Error: %s:%d, %s\n", __FILE__, __LINE__,                   \
                   cudaGetErrorString(error));                                 \
    }

__global__ void stencil_kernel(float *A, int N) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i >= 1 && i < N - 1) {
        float left = A[i - 1];
        float right = A[i + 1];
        A[i] = left + right;
    }
}

int main() {
    const int N = 1 << 20; // 1M元素
    float *h_A = new float[N];
    float *h_result_cuda = new float[N];
    float *d_A;

    // 初始化数据
    for (int i = 0; i < N; i++)
        h_A[i] = 1.0f;

    // 分配设备内存并拷贝数据
    CHECK(cudaMalloc(&d_A, N * sizeof(float)));
    CHECK(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice));

    // ----------------------------
    // 正确性验证（仅执行1次核函数）
    // ----------------------------
    stencil_kernel<<<(N + 255) / 256, 256>>>(d_A, N);
    CHECK(cudaMemcpy(h_result_cuda, d_A, N * sizeof(float),
                     cudaMemcpyDeviceToHost));

    // 验证中间元素是否为2.0，边界保持1.0
    float error_sum = 0;
    for (int i = 1; i < N - 1; i++) {
        error_sum += fabs(h_result_cuda[i] - 2.0f);
    }
    assert(error_sum < 1e-4);
    printf("CUDA验证通过！中间元素值=2.0  %.3f\n", error_sum);

    // ----------------------------
    // 性能测试（每次迭代重置数据）
    // ----------------------------
    const int iterations = 100;
    cudaEvent_t start, stop;
    CHECK(cudaEventCreate(&start));
    CHECK(cudaEventCreate(&stop));

    CHECK(cudaEventRecord(start));
    for (int i = 0; i < iterations; i++) {
        // 每次迭代前重置数据为1.0
        CHECK(cudaMemcpy(d_A, h_A, N * sizeof(float), cudaMemcpyHostToDevice));
        stencil_kernel<<<(N + 255) / 256, 256>>>(d_A, N);
    }
    CHECK(cudaEventRecord(stop));
    CHECK(cudaEventSynchronize(stop));

    float cuda_time_ms;
    CHECK(cudaEventElapsedTime(&cuda_time_ms, start, stop));
    printf("CUDA平均耗时（含数据重置）: %.3f ms\n", cuda_time_ms / iterations);

    // 清理资源
    CHECK(cudaFree(d_A));
    delete[] h_A;
    delete[] h_result_cuda;
    return 0;
}
